Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

change stride to 4 instead of 1 if int8 is not supported #1397

Merged
merged 6 commits into from
Sep 24, 2024
Merged

change stride to 4 instead of 1 if int8 is not supported #1397

merged 6 commits into from
Sep 24, 2024
+55 −3

Conversation

Rekt3421
Copy link
Collaborator

@Rekt3421 Rekt3421 commented Sep 12, 2024 •
edited
Loading

please see issue #1396 for more details

@Rekt3421
Copy link
Collaborator Author

Is a unit test needed for this? Seems a little trivial @rjodinchr ?

rjodinchr
rjodinchr requested changes Sep 12, 2024
View reviewed changes
Copy link
Collaborator

@rjodinchr rjodinchr left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is not the only place where we create spv::DecorationArrayStride.
I would like to have every place fixed.

Also, why 2? when int8 is not supported, it is converted to int32 not int16.

We would definitely need tests here. At least the one from the issue.

@rjodinchr
Copy link
Collaborator

This one feels like a good one as well:

// RUN: clspv -int8=0 %target %s -o %t.spv
// RUN: spirv-dis -o %t2.spvasm %t.spv
// RUN: FileCheck %s < %t2.spvasm
// RUN: spirv-val --target-env vulkan1.0 %t.spv

constant uchar b[4] = {[0]=42, [1]=13, [2]=0, [3]=5};

void kernel __attribute__((reqd_work_group_size(4, 1, 1))) foo(global uchar* a)
{
  *a = b[get_local_id(0)];
}

// CHECK-DAG:  [[uint:%[^ ]+]] = OpTypeInt 32 0
// CHECK-DAG:  OpTypePointer StorageBuffer [[struct:%[^ ]+]]
// CHECK-DAG:  [[struct]] = OpTypeStruct [[runtimearr:%[^ ]+]]
// CHECK-DAG:  OpDecorate [[runtimearr]] ArrayStride 4

You will note that even if you change ArrayStride 4 to ArrayStride 2 in the test, it is still not passing with 2. It would be failing with the following message:

Structure id 15 decorated as Block for variable in StorageBuffer storage class must follow standard storage buffer layout rules: member 0 contains an array with stride 2 not satisfying alignment to 4

@Rekt3421 Rekt3421 changed the title change stride to 2 instead of 1 if int8 is not supported change stride to 4 instead of 1 if int8 is not supported Sep 16, 2024
@Rekt3421 Rekt3421 requested a review from rjodinchr September 16, 2024 15:30
@Rekt3421 Rekt3421 marked this pull request as draft September 16, 2024 15:31
@Rekt3421
Copy link
Collaborator Author

Will add the test from the issue as well.

@Rekt3421 Rekt3421 marked this pull request as ready for review September 16, 2024 15:52
rjodinchr
rjodinchr requested changes Sep 18, 2024
View reviewed changes
lib/SPIRVProducerPass.cpp Outdated Show resolved Hide resolved
@rjodinchr
Copy link
Collaborator

The macos build are failing because this branch has not been rebase on ToT.

@rjodinchr rjodinchr merged commit efc9528 into google:main Sep 24, 2024
10 of 12 checks passed
@Rekt3421 Rekt3421 deleted the 1396 branch September 24, 2024 16:55
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@rjodinchr rjodinchr rjodinchr requested changes

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
3 participants
@Rekt3421 @rjodinchr @kokoro-team